record(GEMMA4-ROCM-KEEP): the 2x R9700 KEEP recipe, and the denominator it still lacks - #676
Conversation
|
Thanks for this — and welcome. Reviewed as part of a sweep over the open external PRs. Before the findings, the part that matters: all eight checkers you list as passing locally do pass — I ran every one at your head SHA and confirmed it. Everything below is something no checker in this repo looks for, so none of it is a diligence failure on your part. Several of your calls are ones this project learned the hard way: disclosing the dropped 1170 first-rep with its cause rather than quietly taking the median, per-depth rep counts, explicitly writing "does not claim that bar", refusing to name an LLVM defect from a spill count, and not touching The blocker is that the recipe cannot be reproduced from Four of the decode knobs are read by no production code in this tree:
Meanwhile every prefill item — SharedK-WMMA on, FLASH off, I don't think the numbers are wrong. The likely story is honest measurements of your working branch written up against the wrong baseline. But a record whose whole value is reproducibility has to name the tree it describes, and the spec header still says Why this was easy to miss, and not your fault: Second blocker: no issue. AGENTS.md wants one linked in three places that agree — the roadmap issue table, the row's spec, and the PR body. The spec header cites On the denominator. The comparison arm is "Vulkan Q8" with the stack unnamed. The prior entry for this same box labelled it Your "does not claim that bar" framing is the right instinct and it is one sentence from being fully admissible. There's a template already in the tree at Placement. The measurement landed on
Two smaller things worth having:
Please keep the residual section as-is. The HIP cm1 339-spill vs ACO 0-spill contrast on the same silicon, with the explicit refusal to call it a named defect, is disciplined work and it should not be lost when the spec is compacted — I'd suggest an issue so it survives independently. Decode regime would help too: 55.5 / 49.1 t/s with no concurrency, depth or batch size can't be compared to anything, including a future vLLM leg. Your prefill depth curve is exactly right; decode just needs the same treatment. Nothing here needs new measurements. Happy to look again once the recipe names its tree. |
localai-org-maint-bot
left a comment
There was a problem hiding this comment.
The measured KEEP result needs the repository-required benchmark checkpoint before this can merge. This change publishes binding throughput numbers in FEATURES/ROCM/USAGE, but it does not update docs/BENCHMARKS.md and provides no committed command/log evidence anchor for reproducing the contributor-lab medians or the discarded outlier. Please add the exact workload/commands and evidence location to BENCHMARKS (or mark the numbers non-binding until that evidence exists), as required by AGENTS.md.
|
Following up on the review above — there is one question only you can answer, and it blocks the rest. The question: which tree produced these numbers?The KEEP recipe lists four decode knobs, and none of them is read by production code on
Meanwhile every prefill item — SharedK-WMMA on, FLASH off, I am not suggesting the measurements are wrong. The most likely explanation is that they are honest numbers from your working branch, written up against So: which revision were the 2112 / 2014 / 1705 / 1099 prefill figures and the 55.5 / 49.1 decode figures taken on? Once that is stated, there are two clean ways forward and either is fine:
The practical urgency is that Not your fault, and worth saying
The smaller items, unchanged from the earlier review
Please keep the residual section as-isThe HIP cm1 339-spill vs ACO 0-spill contrast on the same silicon, with the explicit refusal to call it a named LLVM defect, is the most reusable thing in this PR. I would suggest an issue for it so it survives the spec being compacted later. Nothing here needs new measurements — just the provenance question answered. |
|
This is a legitimate record-only PR — the record is the work here, not a The RADV/ACO comparison is the most valuable result in this branch. Matched On "plateau": in substance you are within the rule, but the title is not. The blocker is not the wording, though — it is that the recipe does not run on
So Relatedly, The measurement is not yet admissible as published. Against The axis that is missing matters more than the word "plateau". Every gap you Two CI notes so they are not re-derived. Your What would make this landable: the revision the numbers came from, the knob |
Contributor closeout (PREFIX_CACHE=0, unique pads, 2026-08-13): 2014 t/s @~11k, 1099 t/s @~42k, decode ~55 t/s. Quality Paris/63/tool_calls. Vulkan Q8 on the same box is still ahead; this is the reliable recipe, not that bar. Speculative/ngram/FMHA/layer-split stay off. HIP cm1 spill is a residual, not a named LLVM defect. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Hermes:grok-4.6 [Hermes]
Matched RADV pipeline stats on gfx1201 (Mesa 26.0.3): 0 spilled VGPR, 0 scratch at VGPR=256 for llama.cpp coopmat1 d=512. HIP cm1 still 235-339 spills. (a)-lean only; no component name, no 3.35x claim. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Hermes:grok-4.6 [Hermes]
…cipe cannot reproduce `docs/ENVIRONMENT.md` described `VT_GEMMA4_PREFILL_GEMM_M` as default `256` over `16..2048`. `gemma4_moe.cpp:1016-1021` accepts `16..8192` and returns `2048`, and has since 2026-08-10. The row was wrong on `main` before this branch touched it and this branch edited that row, so it is repaired in flow. The same edit had dropped a recorded measurement, lab `512` ~+37% prefill vs `64`; AGENTS.md says to move evidence and never to drop it, so it is restored beside the `512`->`2048` ~+80 eng @11k result that explains the default. Four names in the published recipe are read by no product code in this tree. `VT_ATTN_DECODE_KV_SPLITS`, `VT_ATTN_DECODE_SLIDE_SPLITS` and `VT_ATTN_DECODE_SPLIT_WARPS` occur only in `tests/vt/test_gemma4_rocm_fp8_seams.cpp`; `VT_ATTN_DECODE_SLIDE_WARPS` occurs nowhere; `git grep` over `src/` and `include/` returns zero hits for all four. That is mudler#845, whose seam test asserts `EnvInt(name, 16) == 16` with the variable unset and therefore passes whether or not the knob exists. `PEER_ACT` and `PREFILL_GEMM_M` were abbreviated and are now spelled as the product spells them, `VT_GEMMA4_PREFILL_PEER_ACT` and `VT_GEMMA4_PREFILL_GEMM_M`, both already the default. `docs/USAGE.md` therefore publishes only the knobs a reader can set and says so; the spec keeps the full as-run recipe, because a record's job is to hold what happened and a user page's job is to be followable. `docs/FEATURES.md` returns to main's text. A throughput figure's home is `docs/BENCHMARKS.md` under the projection table, this change alters no feature, backend or quantization surface, and fitting the number into a cell already 219 of its 220 characters cost the row its `VT_GEMMA4_*`/`VT_ATTN_*` pointer and the `test_gemma4_rocm_fp8_seams` seam name. `## Owed` names what stays open: no vLLM-ROCm denominator exists for any of these numbers, so `docs/BENCHMARKS.md` keeps this backend at `PENDING: no binding throughput number` and `.agents/benchmark-record.md` is deliberately not appended; the tree SHA behind the 2026-08-13 run is unrecorded; and the Vulkan Q8 gap stays open with the HIP-LLVM versus ACO register-allocation disparity named as the next traceable hypothesis, because AGENTS.md forbids reading it as a ceiling. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [ClaudeCode]
c08af52 to
f649859
Compare
|
@bakon11 — I pushed the fixable parts to The blocker: four knob names in the published recipe are read by nothing in this tree. The first three appear only in This is not you inventing names. It is #845: that seam test asserts So the recipe cannot be followed as written, and I cannot tell which way it fails, because the spec's
Only you know. The SHA, and which of the two, is what unblocks this. I did not guess: the spec keeps your full as-run recipe with that question attached, and What I fixed.
I retitled. The body always said the right thing — "does not close that bar", with the ACO disparity named as the next hypothesis — but on a squash merge only the title reaches The old On the denominator, since it is the other half of why this cannot land as a Every number here is engine-side with nothing on the other side. Under AGENTS.md "Gates" that is not a throughput result, which is why The oracle is not hypothetical — Both items are under |
|
Answering the blocker (coordinator for bakon11's lab, with the lab records in hand): It is your option 1 — the four decode knobs were live on the measured tree, which is not this one.
So: numbers stand as lab measurements with SHA |
…re it exists (#1186) Fixes #785 (taken over from @joral with his blessing). ## Problem `VT_ROCWMMA_OK` is defined only under `__gfx1200__/__gfx1201__`, which exist **only on the device compilation pass** — the host-side `#if defined(VT_ROCWMMA_OK)` guard around both `PagedAttnPrefillSharedKWmma` launches is therefore always false, and the launches are compiled out of the host TU entirely. **The rocWMMA prefill path at d=256/d=512 has never executed**; every such prefill silently falls through to the scalar kernel. ## Fix Host-side **runtime arch dispatch**: query `hipDeviceProp_t::gcnArchName` once per device (prefix match `gfx1200`/`gfx1201`, handles suffixed forms like `gfx1201:xnack-`), launch the WMMA kernel on match, existing scalar path otherwise. Device-side kernel-body guard unchanged, so non-gfx120x fatbin slices still never contain (or launch) WMMA stubs. Decision cached per device — not per build, not per call. ## Evidence (2× R9700, gfx1201, shipping ROCm 7.2.4) Kernel-trace witness on the product seam, A/B on the same build: - **A (gfx1201): exactly one `PagedAttnPrefillSharedKWmma<2,8,16,32,false>` dispatch, no scalar family** — the kernel launches for the first time from an unhacked tree - **B (dispatch forced off): exactly one scalar `PagedAttnPrefillSharedK<2,8,32,32>` dispatch, no WMMA family** - Identical frozen Q/K/V hashes both arms; correctness oracle GREEN; independently reviewed (artifact hashes reproduced exactly) - KD (shipping clang `f58b06d`): 0 spill/private, wave32 ## Expected impact Our lab (which hard-enabled this path locally) measures the KEEP recipe at **~2,014 tok/s @11k / ~1,099 @42k prefill** on 2×R9700 — numbers that were previously unreachable from `main` because the kernel never ran (see the reproducibility discussion on #676). This PR makes them reachable with the published knobs. ## Landing note (maintainer) Rebased from `233c7dc31` onto `main` and force-pushed to the contributor's branch under `maintainerCanModify`, because GitHub reported the pull request as conflicting while `git merge-tree` reported it clean. The cause is that GitHub does not honour the `merge=union` driver that `.gitattributes` sets on `.agents/issue-index.md`, so any branch appending one issue-index row reads as conflicting. The rebase itself produced a defect, which is recorded here because it is a property of the union driver rather than of this change. Replaying turned the branch's in-place update of its `#785` row into a second append, leaving two rows where the branch had one and resurrecting a superseded wording. `check-agent-record.py` refused it. It was repaired at the commit where it arose, keeping the final d=256-only wording, and the result was verified: the `+`/`-` content-line sets of the pull request diff are identical before and after the rebase, authorship is preserved on all five commits, and `check-agent-record`, `check-doc-checkpoint`, `check-issue-index-append-only` and `check-commit-trailers` all pass. `windows-msvc-cpu` and `windows-msvc-vulkan` are red here as they are on every pull request in this repository: those jobs are pull-request-only, have no `main` baseline, and carry no verdict about this change. `commit-protocol-tag` was red because the pull-request body carried no trailer paragraph; this squash body supplies it, and the strict trailer walk over the branch commits reported `OK: commit trailer contract`. On the pre-rebase head every other gate passed, including `build`, both `verify` lanes, both `sanitize-cpu` lanes, `device-leakage`, `agent-record` and `documentation-checkpoint`. The rebase reset those runs and the repository CI queue had not drained them at merge time, so this landed on the pre-rebase verdict plus the local gate run described above. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: Hermes:grok-4.6 [Hermes]
Brings the branch up to `main` so GitHub can compute the merge, and resolves the three record conflicts as unions rather than picks, per AGENTS.md: take the complete target-branch version, then re-apply the scoped edit. `.agents/benchmark-record.md` keeps both appended sections, `main`'s ENG-CUDAGRAPH-BREAK W1 entry and this branch's ROCM-GEMMA4-PREFILL-PEER-BARRIER entry, because it is an append-only log and either side alone drops a measurement. `docs/STATUS.md` keeps `main`'s M4 gate results, which landed with mudler#559, and re-appends this branch's mudler#1047 item-3 attribution sentence. `docs/USAGE.md` keeps `main`'s KEEP recipe and its "This path does" wording from mudler#676, with this branch's prefill-peer lifetime sentence re-applied ahead of it. Every conflicted path was a record or documentation file. No `src/` or `include/` path conflicted, so no product source was hand-resolved here; the product changes this merge carries in are `main`'s own. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude:claude-opus-5 [claude-code]
Brings the branch up to `main` so GitHub can compute the merge, and resolves the two conflicts as unions rather than picks. `tests/CMakeLists.txt` keeps both test registrations, `main`'s `test_gemma4_moe_device_arm_guard` from mudler#1205 and this branch's `test_gemma4_indexed_max_t`, because the registrations are additive and either side alone silently drops a suite. `docs/USAGE.md` needed care and nearly lost content. This branch predates mudler#676, so its paragraph does not carry the contributor KEEP recipe that landed on `main` in `1dac4f9a7`. Taking either side whole was wrong: the branch alone deletes the 2014 t/s recipe, its no-denominator paragraph and the mudler#845 note, while `main` alone drops the `VT_GEMMA4_DECODE_INDEXED_MAX_T` documentation and asserts "No new env", which this branch falsifies by adding one. The resolution keeps the branch's env list and indexed T=2..63 prose, then continues into `main`'s complete text from "This path does" onward, so the KEEP recipe survives intact. Every conflicted path was a build or documentation file. No `src/` or `include/` path conflicted, so no product source was hand-resolved here. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: claude:claude-opus-5 [claude-code]
…and Gemma-4 records `main` gained a ROCm ReshapeAndCache/PagedAttention composition test (#497), the ROCm M4 near-tie gate for Qwen3.5-0.8B GDN (#559), and the Gemma-4 ROCm KEEP recipe record (#676) while this branch was under gate. None of them touches the LTX-2.5 path, the CPU GEMM seam, or the two test files this branch edits. Merged rather than left behind so that the branch's gate runs against what it will land on, and so the trailer gates — which decline to run at all while the branch is behind — have an ancestor to compare against. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [claude-code]
…sh (#1047) fix(rocm/gemma4): #839 retirement-safe prefill peer GeGLU Launch/Finish Fixes #839. ## The defect The Gemma-4 FP8 cross-device prefill helper was one monolithic function. Its donor unpinned the dequant cache immediately after enqueueing the GEMMs and before `hipEventRecord(ev_e)`, and the cache is process-wide, so a second worker could evict a zero-pin slot while the first expert stream was still reading it. Under batched-MoE prefill that is a use-after-free, and the observed symptom is a wedge after matched BEGIN/END with the process idle in `kfd_wait_on_events`. ## What changed `RunGemma4Fp8ExpertGeGLUPrefillOnExpertDevice` stays the public symbol and becomes a wrapper over `LaunchGemma4Fp8ExpertGeGLUPrefillPeer` and `FinishGemma4Fp8ExpertGeGLUPrefillPeer`. The cache pin now lives on the `PeerSlot` and is released only after the host has **observed** retirement, never after a mere enqueue: - Launch pins under the same mutex scope as `GetLocked` and stores `{cache_pin, cache_dev}` on the slot before any GEMM is enqueued. It never unpins. - Finish host-waits `ev_e`, then `hipStreamSynchronize` on the compute stream that received the output copy, and only then unpins under the lock. - A failed fill or ready-event keeps a fill lease; the lease is retired with `RetireFillLocked` after the producer-stream sync and never under `cache.mu`. - A restore failure after publish retires before it rethrows; a failed retire quarantines the pin rather than dropping it. - The same-dev arm follows the same rule. "Enqueue then unpin" is forbidden on both arms, because a `hipStreamWaitEvent` on the compute stream is not a host-side retirement proof. Peer-pipe overlap stays default OFF, so Launch and Finish run back-to-back on slot 0 and the event order is the one shipping today plus slot-scoped scratch. ## What the tests prove `tests/vt/test_ops_gemma4_prefill_peer.cpp` is 23 cases over the host lifetime model and the product source. The lifetime cases drive `HostLaunch`/`HostFinish`; the product is bound to them by source-slice gates that extract the real `Finish` and `RestoreComputeOrThrow` bodies out of `src/vt/rocm/rocm_gemma4_experts.hip`, compile them, run them, and require each mutation to red. Replacing the `hipSetDevice(compute_dev)` throw in the product `RestoreComputeOrThrow` with `(void)compute_dev;` fails three assertions. ## CI repairs in this revision Three checks were red for reasons that had nothing to do with the change under review, and all three are green on `main`. `build-test-cpu` and both `sanitize-cpu` legs failed because the test named one absolute HIP compiler from the contributor's box, `/opt/rocm-7.2.4/core-7.14/bin/hipcc`, in four places and CHECKed that it returned 0 — so any machine without that exact path **failed** rather than reporting that the gate had not run. The toolchain is now resolved (`VLLM_CPP_HIPCC`, `HIPCC`, `ROCM_PATH/bin/hipcc`, `hipcc` on PATH, `/opt/rocm/bin/hipcc`, each probed with `--version`) the way `tests/vt/test_ops_getblas_product.cpp:15-21` resolves its own precondition, and its absence prints a loud `HIP COMPILE GATE NOT RUN` banner on stderr instead of asserting. Measured on this tree: 23/23 cases and 199 assertions with no toolchain plus two banners, 23/23 and 205 assertions with `VLLM_CPP_HIPCC` set, the six recovered assertions being exactly the HIP legs. Both arms still red on the `RestoreComputeOrThrow` mutation, so the repair stopped the gate failing on absent hardware without widening it. Absent hipcc does not `exit(77)` the way `test_ops_getblas_product.cpp` does. That file is one HIP gate and nothing else, while this one carries 21 host-lifetime cases that are the only gate CI has on this change, and exiting would take those with it. The g++ leg of both compile gates still runs, so neither case can report a zero-assertion pass. `pr-size` refused the whole change with `could not classify` because nothing in `scripts/check-pr-size.py` matches `.agents/evidence/`. The five donor slices move beside the spec as `SPEC_EVIDENCE` `.log` files (`.txt` is unclassified too) and the manifest becomes `.agents/specs/rocm-gemma4-prefill-peer-helper-donor.md`. Renaming does not touch bytes and all five SHA256 values in the table verify unchanged. `docs/FEATURES.md` returns to `main`'s text. The branch had rewritten that row and lost the spec link, the `VT_GEMMA4_*`/`VT_ATTN_*` pointer and the `test_gemma4_rocm_fp8_seams` seam name in order to fit; `main`'s cell is already 219 of the 220 characters `check-public-doc-tables` allows. Nothing here owes that page: `check-doc-checkpoint` has keyed `feature_surface` off a changed `REGISTER_VLLM_MODEL` set since #595, and this change touches no file under `src/vllm/model_executor/models/`. The branch was rebuilt by rebase onto `affc2a7fd`, so it carries no untrailered merge commit and all ten original commits keep their authorship. The `docs/USAGE.md` conflict against #837's landed GetBlas text was resolved as a union: both paragraphs survive. ## What this does NOT close Named here and under `## Owed` in `.agents/specs/rocm-gemma4-prefill-peer-helper.md`. Row `ROCM-GEMMA4-PREFILL-PEER-HELPER` owns them and #839 stays open until they are discharged. - **No measurement for the new blocking retirement.** Finish now host-waits `hipEventSynchronize(ev_e)` and `hipStreamSynchronize(cst)` on every call, and the same-dev arm gains one too, where the path was fully asynchronous before. It runs per expert per layer during prefill, so it serialises a pipeline that used to overlap. The wait is required for correctness and correctness comes first, so it stays; what is missing is the number. Prefill throughput before and after on 2x R9700 needs the hardware, and no CI runner here has it. - **`PeerSlot s[2]` with only slot 0 reachable** from the production wrapper. - **The host simulator is a second implementation.** `HostLaunch`/`HostFinish` in `include/vt/rocm/rocm_gemma4_prefill_dequant_cache.h` are hand-written analogues of the product Launch/Finish, ~250 of that header's lines, compiled into every HIP build. Moving them under `tests/` is the fix; it was not done here because only a ROCm box compiles the one translation unit that consumes the header, so the move cannot be verified where this repair was made. ## Known-unrelated CI `windows-msvc-cpu` and `windows-msvc-vulkan` are red on every open PR from a break predating this branch (#503, #584). ## Landing note (maintainer) `main` was merged into this branch before the squash, and three record conflicts were resolved as unions rather than picks, per AGENTS.md: take the complete target-branch version, then re-apply the scoped edit. `.agents/benchmark-record.md` keeps both appended sections, since it is an append-only log and either side alone drops a measurement. `docs/STATUS.md` keeps the M4 gate results that landed with #559 and re-appends this branch's #1047 item-3 attribution sentence. `docs/USAGE.md` keeps the contributor KEEP recipe that landed with #676 and re-applies this branch's prefill-peer lifetime sentence ahead of it. No `src/` or `include/` path conflicted, so no product source was hand-resolved. The resolved tree was gated locally: the 42-test `test_check_gate_commands` suite, plus `check-agent-record`, `check-doc-checkpoint`, `check-issue-index-append-only`, `check-readme-structure`, `check-public-doc-tables` and `check-env-doc`, all green. On the CI verdict, stated precisely. This branch's own product changes are two files, `include/vt/rocm/rocm_gemma4_prefill_dequant_cache.h` and `src/vt/rocm/rocm_gemma4_experts.hip`, and the host lifetime suite `tests/vt/test_ops_gemma4_prefill_peer.cpp` that exercises the state machine. That host suite is what the 17 passing gates on head `de8d4760f` covered. The HIP compilation itself is not built by any lane in this repository's CI, so no CI run here has ever verified it, before or after this merge; the gfx1201 evidence is the contributor's. `windows-msvc-cpu` and `windows-msvc-vulkan` were the only two red gates and are red on every pull request here. The post-merge re-run was still queued behind 19 other jobs at merge time. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: ClaudeCode:claude-opus-5 [ClaudeCode]
record(GEMMA4-ROCM-KEEP): the 2x R9700 KEEP recipe, and the denominator it still lacks
Records the contributor KEEP recipe for Gemma-4-26B-A4B-it FP8 on dual R9700
(gfx1201, ROCm 7.2.4) and what it is still missing. Documents only: no kernel
change, no default change, no env change.
What is recorded
Fair protocol
PREFIX_CACHE=0plus unique pads, 2026-08-13.Decode stream 55.5 t/s at temp=0, 49.1 t/s at temp=0.7. Paris, arith
63andgemma4tool_callsheld. Same-box Vulkan Q8 unique-pad prefill is 3503 @11k and2714 @42k, so the ROCm path is 1.74x and 2.47x behind it there.
The rejected levers are named with their numbers: FMHA_WMMA2 (quality fail plus
0.67x @11k), isolated P1 cm1 wg256 (~1.13x where ~3.35x isolated is needed),
layer-split FIFO (~0.60x), Head-TP peer-read, hipBLASLt dual-GPU. So is the
residual: the faithful HIP cm1 hsaco spills 339 VGPR against RADV/ACO's 0 on an
equivalent llama.cpp coopmat1 spec, which is the mechanism behind the isolated
1.13x and is not yet a named LLVM component.
What is NOT recorded, and why
There is no denominator. Every number above is engine-side with nothing on
the other side of it — no pinned vLLM-ROCm run on the same box, model,
quantization, request shape, concurrency and cache policy. Under AGENTS.md
"Gates" that is not a throughput result, so
docs/BENCHMARKS.mdkeeps thisbackend at
PENDING: no binding throughput numberand this PR does not touch thatrow.
.agents/benchmark-record.mdis also deliberately not appended: it is theappend-only measurement log, and a figure that enters it gets quoted afterwards as
measured.
The oracle is not hypothetical.
docs/ROCM.md§5 documents two working DockervLLM-ROCm recipes on this hardware family, the second building this project's
pinned commit
555967922insiderocm/vllm-dev:basein about 6.5 minutes. Whatis missing is a run, and only the contributor has 2x R9700.
The recipe is not reproducible as written. Four of the names in it are read by
no product code in this tree.
VT_ATTN_DECODE_KV_SPLITS,VT_ATTN_DECODE_SLIDE_SPLITSandVT_ATTN_DECODE_SPLIT_WARPSoccur only intests/vt/test_gemma4_rocm_fp8_seams.cpp;VT_ATTN_DECODE_SLIDE_WARPSoccursnowhere at all;
git grepoversrc/andinclude/returns zero hits for allfour. That is #845, whose seam test asserts
EnvInt(name, 16) == 16with thevariable unset — a tautology that passes whether or not the knob exists, which is
how names nothing reads came to look real.
The
**Status:**line names the run's tree asPR tip feat/gemma4-rocm-fp8-splitwith no commit SHA, so this repository cannot establish which tree produced the
numbers. Either the four decode splits were live on that tree and it is not this
one, or they were inert and the shipped defaults produced these figures. Only
the contributor can say which, and until he does the decode figure has no recipe
behind it. Both items are listed under
## Owedin.agents/specs/gemma4-rocm-fp8-moe.md.Repairs in this revision
docs/ENVIRONMENT.mddescribedVT_GEMMA4_PREFILL_GEMM_Mas default256over16..2048.gemma4_moe.cpp:1016-1021accepts16..8192and returns2048, andhas since 2026-08-10. The row was already wrong on
mainand this branch editsthat exact row, so it is repaired in flow. The same edit had dropped a recorded
measurement — lab
512~+37% prefill vs64— and AGENTS.md says to moveevidence, never to drop it, so it is restored beside the
512→2048~+80 eng@11k result that explains the default.
PEER_ACTandPREFILL_GEMM_Mare now spelled as the product spells them,VT_GEMMA4_PREFILL_PEER_ACTandVT_GEMMA4_PREFILL_GEMM_M, and both are alreadythe default.
docs/USAGE.mdpublishes only knobs a reader can actually set andsays outright that the decode figure is not reproducible from them; the spec keeps
the full as-run recipe, because a record holds what happened and a user page has
to be followable.
docs/FEATURES.mdreturns tomain's text. A throughput figure's home isdocs/BENCHMARKS.mdunder the projection table, this change alters no feature,backend or quantization surface, and fitting the number into a cell already 219 of
its 220 characters cost the row its
VT_GEMMA4_*/VT_ATTN_*pointer and thetest_gemma4_rocm_fp8_seamsseam name.The title no longer says "plateau". The PR body always said the right thing — that
this does not close the Vulkan bar, and it names the ACO disparity as the next
hypothesis — but the title is what survives into
git log, and "plateau" readsthere as the ceiling claim AGENTS.md forbids.
The branch was rebuilt by rebase onto
affc2a7fd, so it carries no untraileredmerge commit and both original commits keep their authorship. The
docs/USAGE.mdconflict against #837's landed GetBlas text was resolved as a union.
Known-unrelated CI
windows-msvc-cpuandwindows-msvc-vulkanare red on every open PR from a breakpredating this branch (#503, #584). The previous
sanitize-cpu (address,undefined)red on this branch was
test_ltx2_video, from a lane 215 commits ahead of the oldbase; the rebase carries it away.
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: ClaudeCode:claude-opus-5 [ClaudeCode]